Skip to content

Handle struct return in FP conversion builtins#3610

Open
vmaksimo wants to merge 1 commit intoKhronosGroup:mainfrom
vmaksimo:fp-conversions-sret
Open

Handle struct return in FP conversion builtins#3610
vmaksimo wants to merge 1 commit intoKhronosGroup:mainfrom
vmaksimo:fp-conversions-sret

Conversation

@vmaksimo
Copy link
Copy Markdown
Contributor

__builtin conversion functions can use sret instead of direct returns (e.g. wrapper type sycl::half). When sret is present, operand 0 is the return pointer, not the source value.

__builtin conversion functions can use `sret` instead of direct returns (e.g. wrapper type `sycl::half`). When sret is present, operand 0 is the return pointer, not the source value.
@vmaksimo vmaksimo marked this pull request as ready for review February 26, 2026 14:54
@MrSidims
Copy link
Copy Markdown
Contributor

MrSidims commented Feb 26, 2026

I was thinking about sycl::half and sycl::bfloat issues one other day. Ideally, imho, clang should be switching to spirv target triple for OpenCL/SYCL/HIP/whatever during compilation with spirv target ABI overwritten not mandating to return structs via sret, diverging it from spir ABI.

It's too much of a task to say: hey, lets change ABI instead of adding a small-little W/A in llvm-spirv, so I won't insist. But may be you should consider using -freg-struct-return clang option for device compilation? @bader do you see it causing potential issues?

@bader
Copy link
Copy Markdown
Contributor

bader commented Feb 28, 2026

I'm out of context of the issue this PR is trying to solve, but from my POV, SPIR-V conversion built-ins should not take/return SYCL data types. I would expect them to work with SYCL class members which represent half/bfloat*/fp* types, which should not require to return user-defined types. Does it make sense?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants